package hpc.exercices.exercice2;

import static org.jocl.CL.CL_CONTEXT_PLATFORM;
import static org.jocl.CL.CL_DEVICE_TYPE_ALL;
import static org.jocl.CL.CL_MEM_COPY_HOST_PTR;
import static org.jocl.CL.CL_MEM_READ_ONLY;
import static org.jocl.CL.CL_MEM_READ_WRITE;
import static org.jocl.CL.clBuildProgram;
import static org.jocl.CL.clCreateBuffer;
import static org.jocl.CL.clCreateCommandQueue;
import static org.jocl.CL.clCreateContext;
import static org.jocl.CL.clCreateKernel;
import static org.jocl.CL.clCreateProgramWithSource;
import static org.jocl.CL.clEnqueueNDRangeKernel;
import static org.jocl.CL.clEnqueueReadBuffer;
import static org.jocl.CL.clEnqueueWriteBuffer;
import static org.jocl.CL.clGetDeviceIDs;
import static org.jocl.CL.clGetPlatformIDs;
import static org.jocl.CL.clReleaseCommandQueue;
import static org.jocl.CL.clReleaseContext;
import static org.jocl.CL.clReleaseKernel;
import static org.jocl.CL.clReleaseMemObject;
import static org.jocl.CL.clReleaseProgram;
import static org.jocl.CL.clSetKernelArg;
import hpc.exercices.Profiler;
import hpc.exercices.Utils;

import java.io.IOException;
import java.nio.charset.Charset;

import org.jocl.CL;
import org.jocl.Pointer;
import org.jocl.Sizeof;
import org.jocl.cl_command_queue;
import org.jocl.cl_context;
import org.jocl.cl_context_properties;
import org.jocl.cl_device_id;
import org.jocl.cl_event;
import org.jocl.cl_kernel;
import org.jocl.cl_mem;
import org.jocl.cl_platform_id;
import org.jocl.cl_program;
import org.junit.Assert;

public class LargeArrayScanWebSolution {

	// the OpenCL context
	private static cl_context context;

	// the OpenCL command queue to which the all work will be dispatched
	private static cl_command_queue commandQueue;

	// kernel program
	private static cl_program program;
	
	// device
	private static cl_device_id device;

	// kernels
	private static cl_kernel bScanKernel;
	private static cl_event bScanKernel_event = new cl_event();

	private static cl_kernel pScanKernel;
	private static cl_event pScanKernel_event = new cl_event();

	private static cl_kernel blockAdditionKernel;
	private static cl_event blockAdditionKernel_event = new cl_event();

	// array elements
	private static int elements = (int) Math.pow(2, 20);

	// block size - splitting array after ....
	private static int blockSize = 512;

	public static void main(String[] args) {

		// 1 million
		int[] array = Utils.randomArray(elements);

		// ////////////////////////////////////////////////////////////
		// SEQUENTIAL //

		int[] result_seq;

		long startSeqScan = System.currentTimeMillis();
		result_seq = SequentialScan.scan(array);
		long endSeqScan = System.currentTimeMillis();
		System.out.println("Duration for SeqScan: "
				+ (endSeqScan - startSeqScan) + " ms");

		// debug
		// System.out.println(Arrays.toString(result_seq));
		// ////////////////////////////////////////////////////////////

		// setUp all CL things
		setUp();
		
		// create Kernels
		createKernels();

		int input[] = array;
		int output[] = new int[array.length];
		
		// calculate needed block size
		blockSize = Math.min(blockSize, elements / 2);
		
		// Calculate number of passes required
		float t = (float) (Math.log((float) elements) / Math
				.log((float) blockSize));
		int pass = (int) t;

		// create buffers
		cl_mem[] outputBuffer = new cl_mem[pass];
		for (int i = 0; i < (int) pass; i++) {
			int size = (int) (elements / Math.pow((float) blockSize, (float) i));
			outputBuffer[i] = clCreateBuffer(context, CL_MEM_READ_WRITE,
					Sizeof.cl_int * size, null, null);
		}

		cl_mem[] blockSumBuffer = new cl_mem[pass];
		for (int i = 0; i < (int) pass; i++) {
			int size = (int) (elements / Math.pow((float) blockSize,
					(float) (i + 1)));
			blockSumBuffer[i] = clCreateBuffer(context, CL_MEM_READ_WRITE,
					Sizeof.cl_int * size, null, null);
		}

		cl_mem inputBuffer = clCreateBuffer(context, CL_MEM_READ_ONLY
				| CL_MEM_COPY_HOST_PTR, Sizeof.cl_int * input.length,
				Pointer.to(input), null);

		int tempLength = (int) (input.length / Math.pow((float) blockSize,
				(float) pass));

		cl_mem tempBuffer = clCreateBuffer(context, CL_MEM_READ_WRITE,
				Sizeof.cl_int * tempLength, null, null);
	

		// Enqueue Write Buffer
		clEnqueueWriteBuffer(commandQueue, inputBuffer, false, 0, input.length
				* Sizeof.cl_int, Pointer.to(input), 0, null, null);

		bScan(input.length, Pointer.to(inputBuffer),
				Pointer.to(outputBuffer[0]), Pointer.to(blockSumBuffer[0]));

		for (int i = 1; i < (int) pass; i++) {
			bScan((int) (input.length / Math.pow((float) blockSize, (float) i)),
					Pointer.to(blockSumBuffer[i - 1]),
					Pointer.to(outputBuffer[i]), Pointer.to(blockSumBuffer[i]));
		}

		// Do scan to tempBuffer
		pScan(tempLength, Pointer.to(blockSumBuffer[pass - 1]),
				Pointer.to(tempBuffer));

		// Do block-addition on outputBuffers
		bAddition((int) (input.length / Math.pow((float) blockSize,
				(float) (pass - 1))), Pointer.to(tempBuffer),
				Pointer.to(outputBuffer[pass - 1]));

		for (int i = pass - 1; i > 0; i--) {
			bAddition((int) (input.length / Math.pow((float) blockSize,
					(float) (i - 1))), Pointer.to(outputBuffer[i]),
					Pointer.to(outputBuffer[i - 1]));

		}

		// result
		clEnqueueReadBuffer(commandQueue, outputBuffer[0], true, 0L,
				(long) (input.length * Sizeof.cl_int), Pointer.to(output), 0,
				null, null);

		// release buffers
		clReleaseMemObject(inputBuffer);
		clReleaseMemObject(tempBuffer);
		for (cl_mem memObject : outputBuffer) {
			clReleaseMemObject(memObject);
		}
		for (cl_mem memObject : blockSumBuffer) {
			clReleaseMemObject(memObject);
		}
		
		// free kernels, context, ....
		shutdown();

		// debug
		// System.out.println(Arrays.toString(output));
		// System.out.println(Arrays.toString(result_seq));

		// sequential result equal parallel result
		Assert.assertArrayEquals(result_seq, output);

	}

	private static void setUp() {
		final int platformIndex = 0;
		final long deviceType = CL_DEVICE_TYPE_ALL;
		final int deviceIndex = 0;

		// enable exceptions
		CL.setExceptionsEnabled(true);

		// obtain the number of platforms
		int numPlatformsArray[] = new int[1];
		clGetPlatformIDs(0, null, numPlatformsArray);
		int numPlatforms = numPlatformsArray[0];

		// obtain a platform ID
		cl_platform_id platforms[] = new cl_platform_id[numPlatforms];
		clGetPlatformIDs(platforms.length, platforms, null);
		cl_platform_id platform = platforms[platformIndex];

		// initialize the context properties
		cl_context_properties contextProperties = new cl_context_properties();
		contextProperties.addProperty(CL_CONTEXT_PLATFORM, platform);

		// obtain the number of devices for the platform
		int numDevicesArray[] = new int[1];
		clGetDeviceIDs(platform, deviceType, 0, null, numDevicesArray);
		int numDevices = numDevicesArray[0];

		// obtain a device ID
		cl_device_id devices[] = new cl_device_id[numDevices];
		clGetDeviceIDs(platform, deviceType, numDevices, devices, null);
		device = devices[deviceIndex];

		// create a context for the selected device
		context = clCreateContext(contextProperties, 1,
				new cl_device_id[] { device }, null, null, null);

		// tried without deprecated method, not successful cl_queue_properties
		/*
		 * cl_queue_properties queueProperties = new cl_queue_properties();
		 * queueProperties.addProperty(CL.CL_QUEUE_PROPERTIES,
		 * CL.CL_QUEUE_PROFILING_ENABLE);
		 * queueProperties.addProperty(CL.CL_QUEUE_PROPERTIES,
		 * CL.CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE);
		 * 
		 * // Create a command-queue for the selected device cl_command_queue
		 * cl_command_queue commandQueue =
		 * CL.clCreateCommandQueueWithProperties(context, device,
		 * queueProperties, null);
		 */
		
		long properties = 0;
		properties |= CL.CL_QUEUE_PROFILING_ENABLE;
		properties |= CL.CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE;

		// create a command-queue for the selected device cl_command_queue
		commandQueue = clCreateCommandQueue(context, device, properties, null);
	}

	private static boolean bScan(int length, Pointer inputBuffer,
			Pointer outputBuffer, Pointer blockSumBuffer) {

		// global worksize (total work items)
		long globalWorkSize[] = new long[] { length / 2 };

		// local worksize
		long localWorkSize[] = new long[] { blockSize / 2 };

		// inputBuffer
		clSetKernelArg(bScanKernel, 0, Sizeof.cl_mem, inputBuffer);

		// outputBuffer
		clSetKernelArg(bScanKernel, 1, Sizeof.cl_mem, outputBuffer);

		// local memory
		clSetKernelArg(bScanKernel, 2, blockSize * Sizeof.cl_int, null);

		// block_size
		clSetKernelArg(bScanKernel, 3, Sizeof.cl_int,
				Pointer.to(new int[] { blockSize }));

		// length
		clSetKernelArg(bScanKernel, 4, Sizeof.cl_int,
				Pointer.to(new int[] { length }));

		// sum of blocks
		clSetKernelArg(bScanKernel, 5, Sizeof.cl_mem, blockSumBuffer);

		// execute the kernel
		clEnqueueNDRangeKernel(commandQueue, bScanKernel, 1, null, globalWorkSize,
				localWorkSize, 0, null, bScanKernel_event);

		CL.clWaitForEvents(1, new cl_event[] { bScanKernel_event });

		System.out.println("bScan Kernel:");
		Profiler.profile(bScanKernel_event,true);

		return true;
	}

	private static boolean pScan(int length, Pointer inputBuffer,
			Pointer outputBuffer) {

		long[] globalWorkSize = new long[] { length / 2 };
		long[] localWorkSize = new long[] { length / 2 };

		// inputBuffer
		clSetKernelArg(pScanKernel, 0, Sizeof.cl_mem, inputBuffer);

		// outputBuffer
		clSetKernelArg(pScanKernel, 1, Sizeof.cl_mem, outputBuffer);

		// local memory
		clSetKernelArg(pScanKernel, 2, length * Sizeof.cl_int, null);

		// length
		clSetKernelArg(pScanKernel, 3, Sizeof.cl_int,
				Pointer.to(new int[] { length }));

		// execute the kernel
		clEnqueueNDRangeKernel(commandQueue, pScanKernel, 1, null,
				globalWorkSize, localWorkSize, 0, null, pScanKernel_event);

		CL.clWaitForEvents(1, new cl_event[] { pScanKernel_event });

		System.out.println("pScan Kernel:");
		Profiler.profile(pScanKernel_event,true);

		return true;
	}

	private static boolean bAddition(int length, Pointer inputBuffer,
			Pointer outputBuffer) {

		long[] globalWorkSize = new long[] { length };
		long[] localWorkSize = new long[] { blockSize };

		// inputBuffer
		clSetKernelArg(blockAdditionKernel, 0, Sizeof.cl_mem, inputBuffer);

		// outputBuffer
		clSetKernelArg(blockAdditionKernel, 1, Sizeof.cl_mem, outputBuffer);

		// Enqueue a kernel run call
		clEnqueueNDRangeKernel(commandQueue, blockAdditionKernel, 1, null,
				globalWorkSize, localWorkSize, 0, null, blockAdditionKernel_event);

		CL.clWaitForEvents(1, new cl_event[] { blockAdditionKernel_event });

		System.out.println("Addition Kernel:");
		Profiler.profile(blockAdditionKernel_event,true);

		return true;
	}

	private static void createKernels() {

		String programSource = null;
		try {
			programSource = Utils.readFile(
					"kernelCode/webSolution_largeArrays.cl",
					Charset.defaultCharset());

		} catch (IOException e) {
			System.err.println("kernel source code couldn't be read! Message:"
					+ e.getMessage());
		}

		// create the program from the source code
		program = clCreateProgramWithSource(context, 1,
				new String[] { programSource }, null, null);

		// build the program
		clBuildProgram(program, 0, null, null, null, null);

		// create the kernel
		bScanKernel = clCreateKernel(program, "ScanLargeArrays", null);
		pScanKernel = clCreateKernel(program, "prefixSum", null);
		blockAdditionKernel = clCreateKernel(program, "blockAddition", null);
	}

	/*
	 * free kernels, context, ....
	 */
	private static void shutdown() {
		clReleaseKernel(bScanKernel);
		clReleaseKernel(pScanKernel);
		clReleaseKernel(blockAdditionKernel);
		clReleaseProgram(program);
		clReleaseCommandQueue(commandQueue);
		clReleaseContext(context);
	}
}
